OpenGL FireRender/CUDA Interop Code
In [1]:
ENV["PATH"]="/usr/local/cuda-7.5/bin:"*ENV["PATH"]
ENV["LD_LIBRARY_PATH"]="/usr/local/cuda-7.5/lib64:/usr/local/cuda/nvvm/lib64:"*get(ENV,"LD_LIBRARY_PATH","")
ENV["AFMODE"] = "CUDA" # this example works with CUDA
using ArrayFire, CUDArt, GLAbstraction, Cxx, ModernGL
type CUDAGLBuffer{T} <: GLAbstraction.GPUArray{T, 1}
buffer::GLBuffer{T}
graphics_resource::Ref{CUDArt.rt.cudaGraphicsResource_t}
ismapped::Bool
end
function CUDAGLBuffer(buffer::GLBuffer, flag = 0)
cuda_resource = Ref{CUDArt.rt.cudaGraphicsResource_t}(C_NULL)
CUDArt.rt.cudaGraphicsGLRegisterBuffer(cuda_resource, buffer.id, flag)
CUDAGLBuffer(buffer, cuda_resource, false)
end
function map_resource(buffer::CUDAGLBuffer)
if !buffer.ismapped
CUDArt.rt.cudaGraphicsMapResources(1, buffer.graphics_resource, C_NULL)
buffer.ismapped = true;
end
nothing
end
function unmap_resource(buffer::CUDAGLBuffer)
if buffer.ismapped
CUDArt.rt.cudaGraphicsUnmapResources(1, buffer.graphics_resource, C_NULL)
buffer.ismapped = false
end
nothing
end
function copy_from_device_pointer{T}(
cuda_mem_ptr::Ptr{T},
cuda_gl_buffer::CUDAGLBuffer,
)
map_resource(cuda_gl_buffer)
buffersize = length(cuda_gl_buffer.buffer)*sizeof(eltype(cuda_gl_buffer.buffer))
if cuda_gl_buffer.buffer.buffertype == GL_RENDERBUFFER
array_ptr = Ref{CUDArt.rt.cudaArray_t}(C_NULL)
CUDArt.rt.cudaGraphicsSubResourceGetMappedArray(array_ptr, cuda_gl_buffer.graphics_resource[], 0, 0)
CUDArt.rt.cudaMemcpyToArray(array_ptr[], 0, 0, cuda_mem_ptr, buffersize, CUDArt.rt.cudaMemcpyDeviceToDevice)
else
opengl_ptr = Ref{Ptr{Void}}(C_NULL); size_ref = Ref{Csize_t}(buffersize)
CUDArt.rt.cudaGraphicsResourceGetMappedPointer(opengl_ptr, size_ref, cuda_gl_buffer.graphics_resource[])
CUDArt.rt.cudaMemcpy(opengl_ptr[], cuda_mem_ptr, buffersize, CUDArt.rt.cudaMemcpyDeviceToDevice)
end
unmap_resource(cuda_gl_buffer)
end
"""
Gets the device pointer from the mapped resource
Sets is_mapped to true
"""
function copy_to_device_pointer{T}(
cuda_mem_ptr::Ptr{T},
cuda_gl_buffer::CUDAGLBuffer,
)
map_resource(cuda_gl_buffer)
is_mapped = true
buffersize = length(cuda_gl_buffer.buffer)*sizeof(eltype(cuda_gl_buffer.buffer))
if cuda_gl_buffer.buffer.buffertype == GL_RENDERBUFFER
array_ptr = Ref{CUDArt.rt.cudaArray_t}(C_NULL);
CUDArt.rt.cudaGraphicsSubResourceGetMappedArray(array_ptr, cuda_gl_buffer.graphics_resource[], 0, 0)
CUDArt.rt.cudaMemcpyFromArray(cuda_mem_ptr, array_ptr[], 0, 0, buffersize, CUDArt.rt.cudaMemcpyDeviceToDevice)
else
opengl_ptr = Ref{Ptr{Void}}(C_NULL); size_ref = Ref{Csize_t}(buffersize)
CUDArt.rt.cudaGraphicsResourceGetMappedPointer(opengl_ptr, size_ref, cuda_gl_buffer.graphics_resource[])
CUDArt.rt.cudaMemcpy(cuda_mem_ptr, opengl_ptr, buffersize, CUDArt.rt.cudaMemcpyDeviceToDevice)
end
unmap_resource(cuda_gl_buffer)
end
# ArrayFire.AFArray
function Base.copy(source::ArrayFire.AFArray, target::CUDAGLBuffer)
d_ptr = ArrayFire.af_device(source)
copy_from_device_pointer(d_ptr, target)
end
function Base.copy(source::CUDAGLBuffer, target::ArrayFire.AFArray)
d_ptr = ArrayFire.af_device(target)
copy_to_device_pointer(d_ptr, target)
end
GLVisualize Setup
In [2]:
using GLVisualize, GeometryTypes, Colors
w=glscreen()
nothing
Simulation
In [4]:
width = 5f0; height = 5f0;
function simulate(pos, vels, forces, dt)
pos[1] += vels[1] * dt
pos[2] += vels[2] * dt
pos[3] += vels[3] * dt
#calculate distance to center
diff_x = pos[1] - width/2
diff_y = pos[2] - height/2
diff_z = pos[3] - height/2
dist = sqrt( diff_x.*diff_x + diff_y.*diff_y + diff_z.*diff_z)
#calculate normalised force vectors
forces[1] = -1 * diff_x ./ dist;
forces[2] = -1 * diff_y ./ dist;
forces[3] = -1 * diff_z ./ dist;
#update force scaled to time and magnitude constant
forces[1] *= dt
forces[2] *= dt
forces[3] *= dt
#dampening
vels[1] *= 1 - (0.0005*dt);
vels[2] *= 1 - (0.0005*dt);
vels[3] *= 1 - (0.0005*dt);
#update velocities from forces
vels[1] = vels[1] .+ forces[1];
vels[2] = vels[2] .+ forces[2];
vels[3] = vels[3] .+ forces[3];
end
afmax(a,b) = AFArray{Float32}(ArrayFire.af_max(a, b))
afmin(a,b) = AFArray{Float32}(ArrayFire.af_min(a, b))
function collisions(pos, vels)
#clamp particles inside screen border
projected_px = afmin(width, afmax(0, pos[1]))
projected_py = afmin(height - 1, afmax(0, pos[2]))
projected_pz = afmin(height - 1, afmax(0, pos[3]))
#calculate distance to center
diff_x = projected_px - width/2
diff_y = projected_py - height/2
diff_z = projected_pz - height/2
dist = sqrt( diff_x.*diff_x + diff_y.*diff_y + diff_z.*diff_z)
#collide with center sphere
radius = 0.01;
elastic_constant = 0.91f0
dr = dist.<radius
s = icxx"af::sum<int>($dr);"
if s > 0
vels[1][dr] = -elastic_constant * vels[1][dr]
vels[2][dr] = -elastic_constant * vels[2][dr]
vels[3][dr] = -elastic_constant * vels[3][dr]
#normalize diff vector
diff_x = diff_x ./ dist
diff_y = diff_y ./ dist
diff_z = diff_z ./ dist
#place all particle colliding with sphere on surface
pos[1][dr] = width/2 + diff_x[dr] * radius
pos[2][dr] = height/2 + diff_y[dr] * radius
pos[3][dr] = height/2 + diff_z[dr] * radius
end
end
Out[4]:
Interactive Part
In [5]:
total_particles = 1000;
# Initialize the kernel array just once
# Generate a random starting state
pos = Any[
rand(AFArray{Float32}, total_particles) * width,
rand(AFArray{Float32}, total_particles) * height,
rand(AFArray{Float32}, total_particles) * height,
]
vels = Any[
randn(AFArray{Float32}, total_particles),
randn(AFArray{Float32}, total_particles),
randn(AFArray{Float32}, total_particles),
]
forces = Any[
randn(AFArray{Float32}, total_particles),
randn(AFArray{Float32}, total_particles),
randn(AFArray{Float32}, total_particles),
]
j_pos = ntuple(3) do i
Array(pos[i])
end
intensity = sqrt(forces[1].*forces[1] + forces[2].*forces[2] + forces[3].*forces[3])
nothing
In [6]:
i_julia = Array(intensity)
view(visualize(
(Sphere(Point2f0(0), 0.01f0), j_pos),
billboard=true,
color=nothing,
intensity=i_julia,
color_norm = Vec2f0(0, 0.2),
color_map = map(RGBA{Float32}, colormap("RdBu", 20)),
model = translationmatrix(-Vec3f0(2.5))
), w, camera=:perspective)
nothing
In [7]:
GLWindow.render_frame(w)
gpu = w.renderlist[1][1]
gpu_pos = gpu[:position_x], gpu[:position_y], gpu[:position_z]
gpu_cu_pos = map(gpu_pos) do buff
CUDAGLBuffer(buff)
end
intensity_cu = CUDAGLBuffer(gpu[:intensity])
nothing
In [24]:
gpu[:color_norm] = Vec2f0(0,0.2)
Out[24]:
In [9]:
tic();
for i=1:1000
dt = toq();tic()
#check for collisions and adjust positions/velocities accordingly
collisions(pos, vels);
#run force simulation and update particles
simulate(pos, vels, forces, dt)
# make device to device copy
map(zip(gpu_cu_pos, pos)) do gcu_af
gcu, af = gcu_af
copy(af, gcu)
end
GLWindow.render_frame(w)
GLWindow.pollevents()
end